#ifndef __MEMORY_H__
#define __MEMORY_H__

#include <cassert>
#include "global.h"

// Borrowed from Minisat 1.14, modified for CUDA device.
// The original Minisat allocation size is "sizeof(T) * size", we use raw size in bytes here.
// The exact raw size number of bytes should be calculated before calling these functions.

#ifdef __CUDA_ARCH__
//=================================================================================================
// 'malloc()'-style memory allocation -- it returns NULL and continues.

// Like the system malloc(), the size is in bytes.
template<class T>
__device__ __forceinline__
T* xmalloc( size_t size )
{
	// Calls CUDA device's malloc. The returned pointer is guaranteed to be aligned to a 16-byte boundary.
	T* tmp = (T*)malloc( size );
#ifdef DEBUG_MEMORY
	printf( "[memory] (device) xmalloc(): %x = malloc(%d)\n", tmp, size );
#endif
	if ( (tmp == NULL) && (size != 0) )
	{ printf( "[memory] (device) xmalloc(): ERROR: malloc() returns NULL.\n" ); }

	return tmp;
}
#else
template<class T> static inline
T* xmalloc(size_t size)
{
    T*   tmp = (T*)malloc(size);
    assert(size == 0 || tmp != NULL);
    return tmp;
}
#endif

#ifdef __CUDA_ARCH__
template<class T>
__device__ __forceinline__
void xfree( T* ptr )
{
	if ( ptr != NULL )
	{
		free((void*)ptr);
#ifdef DEBUG_MEMORY
		printf( "[memory] (device) xfree(%x)\n", ptr );
#endif
	}
	else
	{ printf( "[memory] (device) xfree() receives NULL.\n" ); }
}
#else
template<class T> static inline
void xfree(T *ptr)
{
    if (ptr != NULL) free((void*)ptr);
}
#endif

#ifdef __CUDA_ARCH__
// CUDA device has no realloc(). Just use free() and malloc().
// All malloc() is aligned to a 16-byte boundary.
// It seems this boundary on GTX 480 is 80 bytes.
template<class T>
__device__ __forceinline__
T* xrealloc( T* ptr, size_t oldSize, size_t newSize ) {
	if ( ptr != NULL )
	{
		// To prevent boundary from overflow.
		size_t oldBoundary = ( oldSize == 0 ) ? 16 : ((((oldSize - 1) >> 4) + 1) << 4);
		size_t newBoundary = ( newSize == 0 ) ? 16 : ((((newSize - 1) >> 4) + 1) << 4);
#ifdef DEBUG_MEMORY
		printf( "[memory] (device) xrealloc(): oldBoundary = %d, newBoundary = %d\n", oldBoundary, newBoundary );
#endif
		if ( newBoundary <= oldBoundary )
		{ return ptr; }
	}

	T* tmp = (T*)xmalloc<T>( newSize );
#ifdef DEBUG_MEMORY
	printf( "[memory] (device) xrealloc(): ptr = %x, oldSize = %d, %x = xmalloc(%d)\n", ptr, oldSize, tmp, newSize );
#endif
	if ( (tmp == NULL) && (newSize != 0) )
	{ printf( "[memory] (device) xrealloc() error: malloc() returns NULL.\n" ); }

	if ( ptr != NULL )
	{
		// Copy the exact data in bytes. Copy constructor is not invoked here.
		char* ptrCopy = (char*)ptr;
		char* tmpCopy = (char*)tmp;

		for ( int i = 0; i < oldSize; ++i )
		{ tmpCopy[i] = ptrCopy[i]; }
#ifdef ASSERT
		// Compare each element in higher level.
		for ( int i = 0; i < (oldSize / sizeof(T)); ++i )
		{
			if ( tmp[i] != ptr[i] )
			{ printf( "[memory] (device) xrealloc(): Assertion failed: inconsistent realloced memory content.\n" ); }
		}
#endif
		xfree<T>( ptr );
#ifdef DEBUG_MEMORY
		printf( "[memory] (device) xrealloc(): xfree(%x)\n", ptr );
#endif
	}
	return tmp;
}
#else
template<class T> static inline
T* xrealloc(T* ptr, size_t size)
{
    T*   tmp = (T*)realloc((void*)ptr, size);
    assert(size == 0 || tmp != NULL);
    return tmp;
}
#endif

#endif /* __MEMORY_H__ */
